#ifndef __REDUCE_H__
#define __REDUCE_H__

#include "..\common\framework.h"

template<unsigned int blockSize>
__device__ void warpSumReduce(volatile float* temps, unsigned int tid) {
	if(blockSize >= 64) temps[tid] += temps[tid+32];
	if(blockSize >= 32) temps[tid] += temps[tid+16];
	if(blockSize >= 16) temps[tid] += temps[tid+ 8];
	if(blockSize >=  8) temps[tid] += temps[tid+ 4];
	if(blockSize >=  4) temps[tid] += temps[tid+ 2];
	if(blockSize >=  2) temps[tid] += temps[tid+ 1];
}

template<unsigned int blockSize>
__device__ void sumReduce(float* temps, unsigned int tid) {
	if(blockSize >= 512) { if(tid < 256) {temps[tid] += temps[tid+256];}__syncthreads();}
	if(blockSize >= 256) { if(tid < 128) {temps[tid] += temps[tid+128];}__syncthreads();}
	if(blockSize >= 128) { if(tid <  64) {temps[tid] += temps[tid+ 64];}__syncthreads();}
	if(tid < 32) warpSumReduce<blockSize>(temps,tid);
}

template<unsigned int blockSize>
__device__ void warpMaxReduce(volatile float* temps, unsigned int tid) {
	if(blockSize >= 64) temps[tid] = fmaxf(temps[tid], temps[tid+32]);
	if(blockSize >= 32) temps[tid] = fmaxf(temps[tid], temps[tid+16]);
	if(blockSize >= 16) temps[tid] = fmaxf(temps[tid], temps[tid+ 8]);
	if(blockSize >=  8) temps[tid] = fmaxf(temps[tid], temps[tid+ 4]);
	if(blockSize >=  4) temps[tid] = fmaxf(temps[tid], temps[tid+ 2]);
	if(blockSize >=  2) temps[tid] = fmaxf(temps[tid], temps[tid+ 1]);
}

template<unsigned int blockSize>
__device__ void maxReduce(float* temps, unsigned int tid) {
	if(blockSize >= 512) { if(tid < 256) {temps[tid] = fmaxf(temps[tid], temps[tid+256]);}__syncthreads();}
	if(blockSize >= 256) { if(tid < 128) {temps[tid] = fmaxf(temps[tid], temps[tid+128]);}__syncthreads();}
	if(blockSize >= 128) { if(tid <  64) {temps[tid] = fmaxf(temps[tid], temps[tid+ 64]);}__syncthreads();}
	if(tid < 32)
		warpMaxReduce<blockSize>(temps, tid);
}

template<unsigned int blockSize>
__device__ void warpArgMinReduce(volatile float * values, volatile int * indices, unsigned int tid) {
	if(blockSize >= 64)  if(values[tid + 32] < values[tid]) { values[tid] = values[tid+32];indices[tid] = indices[tid+32]; }
	if(blockSize >= 32)  if(values[tid + 16] < values[tid]) { values[tid] = values[tid+16];indices[tid] = indices[tid+16]; }
	if(blockSize >= 16)  if(values[tid +  8] < values[tid]) { values[tid] = values[tid+ 8];indices[tid] = indices[tid+ 8]; }
	if(blockSize >=  8)  if(values[tid +  4] < values[tid]) { values[tid] = values[tid+ 4];indices[tid] = indices[tid+ 4]; }
	if(blockSize >=  4)  if(values[tid +  2] < values[tid]) { values[tid] = values[tid+ 2];indices[tid] = indices[tid+ 2]; }
	if(blockSize >=  2)  if(values[tid +  1] < values[tid]) { values[tid] = values[tid+ 1];indices[tid] = indices[tid+ 1]; }
}

template<unsigned int blockSize>
__device__ void argminReduce(float* values, int* indices, unsigned int tid) {
	if(blockSize >= 512) { if(tid < 256) { if(values[tid + 256] < values[tid]) { values[tid] = values[tid+256];indices[tid] = indices[tid+256]; }}	__syncthreads(); }
	if(blockSize >= 256) { if(tid < 128) { if(values[tid + 128] < values[tid]) { values[tid] = values[tid+128];indices[tid] = indices[tid+128]; }}	__syncthreads(); }
	if(blockSize >= 128) { if(tid <  64) { if(values[tid +  64] < values[tid]) { values[tid] = values[tid+ 64];indices[tid] = indices[tid+ 64]; }}	__syncthreads(); }
	if(tid < 32)
		warpArgMinReduce<blockSize>(values, indices, tid);
}

template<unsigned int blockSize>
__device__ void warpArgMaxReduce(volatile float * values, volatile int * indices, unsigned int tid) {
	if(blockSize >= 64)  if(values[tid] < values[tid + 32]) { values[tid] = values[tid+32];indices[tid] = indices[tid+32]; }
	if(blockSize >= 32)  if(values[tid] < values[tid + 16]) { values[tid] = values[tid+16];indices[tid] = indices[tid+16]; }
	if(blockSize >= 16)  if(values[tid] < values[tid +  8]) { values[tid] = values[tid+ 8];indices[tid] = indices[tid+ 8]; }
	if(blockSize >=  8)  if(values[tid] < values[tid +  4]) { values[tid] = values[tid+ 4];indices[tid] = indices[tid+ 4]; }
	if(blockSize >=  4)  if(values[tid] < values[tid +  2]) { values[tid] = values[tid+ 2];indices[tid] = indices[tid+ 2]; }
	if(blockSize >=  2)  if(values[tid] < values[tid +  1]) { values[tid] = values[tid+ 1];indices[tid] = indices[tid+ 1]; }
}

template<unsigned int blockSize>
__device__ void argmaxReduce(float* values, int* indices, unsigned int tid) {
	if(blockSize >= 512) { if(tid < 256) { if(values[tid] < values[tid + 256]) { values[tid] = values[tid+256];indices[tid] = indices[tid+256]; }}	__syncthreads(); }
	if(blockSize >= 256) { if(tid < 128) { if(values[tid] < values[tid + 128]) { values[tid] = values[tid+128];indices[tid] = indices[tid+128]; }}	__syncthreads(); }
	if(blockSize >= 128) { if(tid <  64) { if(values[tid] < values[tid +  64]) { values[tid] = values[tid+ 64];indices[tid] = indices[tid+ 64]; }}	__syncthreads(); }
	if(tid < 32)
		warpArgMaxReduce<blockSize>(values, indices, tid);
}

#endif
